#ifndef _CDUD_UTILS_CUH__
#define _CDUD_UTILS_CUH__

#define MIN(a,b) ((a)<(b)?(a):(b))
#define MAX(a,b) ((a)>(b)?(a):(b))
#define HASH(v) (-v-2)
#define CEIL(x,b) (((x)+(b)-1)/(b))
#define ALIGN(x,b) ( (x)%(b)==0?(x):((x)+(b)-(x)%(b)) )
#define ALIGN_DOWN(x, b) ( (x)&(-(b)) )
#define ALIGN_UP(x, b) ( ((x)&(-(b)))+(b) )

#define DISPACTH_SM_TYPE(T)\
template<>\
struct SharedMem<T>\
{ __device__ T& operator[](int indx) { extern __shared__ T sh_##T[]; return sh_##T[indx]; } \
  __device__ T* ptr() { extern __shared__ T sh_##T[]; return sh_##T; } }

#define SH_SKEW(idx) (idx)//((idx>>5)+idx)

template<typename T>
struct SharedMem {};

DISPACTH_SM_TYPE(float);
DISPACTH_SM_TYPE(double);
DISPACTH_SM_TYPE(int);
DISPACTH_SM_TYPE(long);
DISPACTH_SM_TYPE(char);

#define CLEAN_REG(buf,len) \
  for(int i=0; i<len; ++i) buf[i] = 0

template <typename T, int folds>
__device__ __forceinline__ void
shuffle_reduce(uint coop_sz, T &val, int delta, T(*udf)(T*,T))
{
  udf(&val, __shfl_down_sync(0xffffffff, val, delta, coop_sz));
  shuffle_reduce<T, folds/2>(coop_sz, val, delta/2, udf);
}

#define SPEC_SHFL_REDUCE(T) template<> __device__ __forceinline__ void \
shuffle_reduce<T, 1>(uint coop_sz, T &val, int delta, T(*udf)(T*,T)) {}

SPEC_SHFL_REDUCE(float)
SPEC_SHFL_REDUCE(double)

#endif